
__kernel void prefixSumB(__global uint * gin, __global uint * gout, __global uint * gworkgroupResult,__local uint * scratch ){
    int global_idx = get_global_id(0);
    int local_idx = get_local_id(0);  // 0 - 512
    int group_idx = get_group_id(0);
    int ws = 1024;
    gin += 2 * ws * group_idx;
    gout += ws * group_idx;

    if (local_idx == 0){
        scratch[ 1 ] = *(gin + 1);
        scratch[ 0 ] = *(gin);
    }

    barrier(CLK_LOCAL_MEM_FENCE);

    for(int i = 2; i <= 512; i <<= 1){
        if ( local_idx < i ){
            int pos = i + local_idx;
            int res;
            if (pos & 1){
                res = scratch[(pos - 1) / 2];
            }else if ( !(pos & (pos - 1)) ){
                //power of two
                res = gin[pos];
            }else{
                res = gin[pos] + scratch[pos / 2 - 1];
            }
            scratch[ pos ] = res;
            //scratch[ 2 * (i + local_idx)] + scratch[ 2 * (i + local_idx) + 1];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    uint toAdd = gworkgroupResult[group_idx];

    //finish last iteration in two steps
    for(int offset = 1024; offset < 2048; offset += 512){
        int pos = offset + local_idx;
        int res;
        if (pos & 1){
            res = scratch[(pos - 1) / 2];
        }else if ( !(pos & (pos - 1)) ){
            //power of two
            res = gin[pos];
        }else{
            res = gin[pos] + scratch[pos / 2 - 1];
        }
        scratch[ pos ] = toAdd + res;
    }

 //   gout[ local_idx ] = scratch [ local_idx ];
 //   gout[ 512 + local_idx ] = scratch [ 512 + local_idx ];
 //   gout[ 1024 + local_idx ] = scratch [ 1024 + local_idx ];
 //   gout[ 1024 + 512 + local_idx ] = scratch [ 1024 + 512 + local_idx ];

      gout[ local_idx ] = scratch [ 1024 + local_idx ];
      gout[ 512 + local_idx ] = scratch [ 1024 + 512 + local_idx ];

//    if (local_idx == 0 ){
//        gout[0] = scratch[1];
//    }
}
